<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<title>Mali OpenCL SDK v1.1.0: SGEMM</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="navtree.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="resize.js"></script>
<script type="text/javascript" src="navtree.js"></script>
<script type="text/javascript">
  $(document).ready(initResizable);
</script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td id="projectlogo"><img alt="Logo" src="logo.png"/></td>
  <td style="padding-left: 0.5em;">
   <div id="projectname">Mali OpenCL SDK v1.1.0
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.2 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Home</span></a></li>
      <li class="current"><a href="pages.html"><span>Help&#160;and&#160;Tutorials</span></a></li>
      <li><a href="files.html"><span>Files</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
</div><!-- top -->
<div id="side-nav" class="ui-resizable side-nav-resizable">
  <div id="nav-tree">
    <div id="nav-tree-contents">
      <div id="nav-sync" class="sync"></div>
    </div>
  </div>
  <div id="splitbar" style="-moz-user-select:none;" 
       class="ui-resizable-handle">
  </div>
</div>
<script type="text/javascript">
$(document).ready(function(){initNavTree('sgemm_tutorial.html','');});
</script>
<div id="doc-content">
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Macros</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Pages</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div class="header">
  <div class="headertitle">
<div class="title">SGEMM </div>  </div>
</div><!--header-->
<div class="contents">
<div class="textblock"><p>SGEMM (Single-Precision General Matrix Multiplication) OpenCL sample.</p>
<h1><a class="anchor" id="sgemmExampleResult"></a>
Example Result</h1>
<p>By adding these two lines at the beginning of <em>sgemmInitialize</em> function in <a class="el" href="sgemm_8cpp.html">sgemm.cpp</a>, we can generate the same random numbers as in this example:</p>
<div class="fragment"><div class="line"><span class="keyword">const</span> <span class="keywordtype">unsigned</span> seed = 11;</div>
<div class="line">srand(seed);</div>
</div><!-- fragment --><p>Input matrix A: </p>
<pre class="fragment">[0.852691    0.004421  -0.103067 -0.191788]
[-0.23658    0.0336409  0.15781   0.582199]
[-0.0814268 -0.857794  -0.63804  -0.0184786]
[0.793476    0.459307   0.955647 -0.306809]
</pre><p>Input matrix B: </p>
<pre class="fragment">[0.0529994  0.507535 -0.55821  -0.849519]
[-0.929501  0.914186  0.464341 -0.652125]
[0.409218  -0.125776 -0.273086  0.731335]
[-0.371732  0.43648  -0.8001    0.233541]
</pre><p>Input matrix C: </p>
<pre class="fragment">[-0.380438 -0.188046  0.665832  -0.503661]
[-0.262456 -0.278552 -0.5179    -0.965873]
[0.459781   0.720241 -0.22676   -0.719225]
[-0.277435 -0.126954 -0.0564545 -0.142268]
</pre><p>Output matrix C: </p>
<pre class="fragment">[0.0321557  0.347259 -0.225749 -0.897793]
[-0.221897  0.117096 -0.413021  0.333833]
[0.584754  -0.681301 -0.186507  0.0857027]
[0.0925018  0.5558   -0.250792 -0.360579]
</pre><h1><a class="anchor" id="sgemmAlgorithm"></a>
The Algorithm</h1>
<p>This sample performs one matrix multiplication and one matrix addition using single precision computation in the expression:</p>
<p>C = &alpha;AB + &beta;C</p>
<p>Which is written in terms of matrix components as:</p>
<p>C<sub>ij</sub> = &alpha;&sum;<sub>k</sub> A<sub>ik</sub>B<sub>kj</sub> + &beta;C<sub>ij</sub></p>
<p>Where A,B and C are matrices of size 1024 x 1024 and &alpha; and &beta; are scalar constants.</p>
<h1><a class="anchor" id="sgemmImplementation"></a>
Implementation</h1>
<p>Unless otherwise noted, all code snippets come from the OpenCL kernel found in <a class="el" href="sgemm_8cl.html">sgemm.cl</a>.</p>
<p>This kernel vectorizes the summation of matrices A B and C and additionally computes a vector of elements in the resulting matrix C following the equation: </p>
<pre class="fragment">matrixC = alpha * (matrixA * matrixB) + beta * matrixC
</pre><ol type="1">
<li><p class="startli"><b>Choosing the size of the kernel</b></p>
<p class="startli">As we are working with symmetric matrices, we adjust the <em>globalWorksize</em> to be the size of the matrices, so each work-item works on a full row/column. </p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Each kernel outputs one element in matrixC,</span></div>
<div class="line"><span class="comment">     * therefore the total number of work items must be the number of elements (matrixOrder * matrixOrder).</span></div>
<div class="line"><span class="comment">     * To accomplish this we use a global worksize split into 2 dimensions both of matrixOrder size.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    <span class="keywordtype">size_t</span> globalWorksize[2] = {matrixOrder, matrixOrder};</div>
</div><!-- fragment --><p> Then the data is accessed by retrieving the pointer to the row/column with <em>get_global_id</em>. </p>
<div class="fragment"><div class="line">    <span class="keyword">const</span> <span class="keywordtype">int</span> i = get_global_id(1);</div>
<div class="line">    <span class="keyword">const</span> <span class="keywordtype">int</span> j = get_global_id(0);</div>
<div class="line">    float4 sum = (float4)0.0f;</div>
<div class="line">    float4 matrixBColumn;</div>
</div><!-- fragment --></li>
<li><p class="startli"><b>Loading the input data and doing the calculation</b></p>
<p class="startli">Based on the previous description, each work-item will work on one row of matrixA and one column of matrixB to compute the value of one element in the result matrixC.</p>
<p class="startli">For the first part of the <em>for</em> loop, values from a column of matrixB are loaded in sets of 4 floats: </p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Load 4 values from a column of data from matrixB, and 4 values from a row in matrixA,</span></div>
<div class="line"><span class="comment">     * then multiply them together. Repeat until all values in the column/row have been multiplied.</span></div>
<div class="line"><span class="comment">     * We only want the sum of the calculation so we can add the result of each calculation to the last.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> k = 0; k &lt; matrixOrder; k+=4)</div>
<div class="line">    {</div>
<div class="line">        matrixBColumn.x = matrixB[bOffset];</div>
<div class="line">        bOffset += matrixOrder;</div>
<div class="line"></div>
<div class="line">        matrixBColumn.y = matrixB[bOffset];</div>
<div class="line">        bOffset += matrixOrder;</div>
<div class="line"></div>
<div class="line">        matrixBColumn.z = matrixB[bOffset];</div>
<div class="line">        bOffset += matrixOrder;</div>
<div class="line"></div>
<div class="line">        matrixBColumn.w = matrixB[bOffset];</div>
<div class="line">        bOffset += matrixOrder;</div>
</div><!-- fragment --><p> The <em>sum</em> variable accumulates the multiplication of the 4 values of a column from matrixB that have been loaded before, and the next 4 values in a row from matrixA, which are loaded during the multiplication.</p>
<p class="startli">MatrixA pointer is moved to the next 4 values once the loop finish an iteration. </p>
<div class="fragment"><div class="line">        sum += vload4 (0, matrixA) * matrixBColumn;</div>
<div class="line">        matrixA += 4;</div>
<div class="line">    }</div>
</div><!-- fragment --><p> The <em>for</em> statement will loop until completing a row in matrixA and a column in matrixB.</p>
<p class="startli">We are using vector types in the kernel because Mali-T600 series GPUs have 128-bit vector registers and can do arithmetic on vector types. Therefore, we use OpenCL vectors to make more efficient use of the hardware, leading to higher performance.</p>
</li>
<li><p class="startli"><b>Storing the result</b></p>
<p class="startli">Lastly we do the final calculation and store the result in matrixC: </p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Sum the 4 results to get the single output of multiplying a row of matrixA by a column of matrixB.</span></div>
<div class="line"><span class="comment">     * Then carry out the final calculation.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    matrixC[i * matrixOrder + j] = alpha * (sum.x + sum.y + sum.z + sum.w) + beta * matrixC[i * matrixOrder + j];</div>
</div><!-- fragment --> </li>
</ol>
<h1><a class="anchor" id="sgemmRunning"></a>
Running the Sample</h1>
<ol type="1">
<li><p class="startli">From a command prompt in the root of the SDK, run:</p>
<div class="fragment"><div class="line">cd samples/<a class="code" href="sgemm_8cl.html#abfe7be25d42e6bb53f76001bdbab6b25" title="SGEMM kernel function.">sgemm</a></div>
</div><!-- fragment --> <div class="fragment"><div class="line">make install</div>
</div><!-- fragment --><p class="startli">This compiles the SGEMM sample code and copies all the files it needs to run to the bin folder in the root directory of the SDK.</p>
</li>
<li>Copy this folder to the board.</li>
<li><p class="startli">Navigate to the folder on the board and run the SGEMM binary:</p>
<div class="fragment"><div class="line">./<a class="code" href="sgemm_8cl.html#abfe7be25d42e6bb53f76001bdbab6b25" title="SGEMM kernel function.">sgemm</a></div>
</div><!-- fragment --></li>
<li><p class="startli">You should see output similar to:</p>
<div class="fragment"><div class="line">Profiling information:</div>
<div class="line">Queued time:    0.094ms</div>
<div class="line">Wait time:      0.074073ms</div>
<div class="line">Run time:       1316.41ms</div>
</div><!-- fragment --></li>
</ol>
<p>Find solutions for <a class="el" href="common_issues.html">Common Issues</a>.</p>
<h1><a class="anchor" id="sgemmMoreInformation"></a>
More Information</h1>
<p>For more information have a look at the code in <a class="el" href="sgemm_8cpp.html">sgemm.cpp</a> and <a class="el" href="sgemm_8cl.html">sgemm.cl</a>. </p>
</div></div><!-- contents -->
</div><!-- doc-content -->
<!-- start footer part -->
<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
  <ul>
    <li class="navelem"><a class="el" href="tutorials.html">Tutorials</a></li>
    <li class="footer">
        <a href="http://www.arm.com/">(C) ARM Ltd. 2013</a>
    </li>
  </ul>
</div>
</body>
</html>
